#include <iostream>
#include <cuda_runtime.h>
#include "../common/common.h"
#define DEVICENUM 0
#define BDIMX 32 // 方形共享内存行数
#define BDIMY 32 // 方形共享内存列数
#define IPAD 1

__global__ void setRowReadColDyn(int *out)
{
    extern __shared__ int tile[]; // 动态共享内存的申请需用extern关键字
    // 由于设置的block是二维的，但动态共享内存是一维的
    // 所以，基于二维线程索引来计算一维共享内存访问索引方式如下
    // row_idx：根据二维线程索引计算出的一维行主序内存偏移量，即按行展开
    unsigned int row_idx = threadIdx.x + threadIdx.y * blockDim.x;
    // col_idx：根据二维线程索引计算出的一维列主序内存偏移量，即按列展开
    unsigned int col_idx = threadIdx.y + threadIdx.x * blockDim.y;

    tile[row_idx] = row_idx; // 按行主序设置，无冲突
    __syncthreads();         // 等待所有线程完成

    out[row_idx] = tile[col_idx]; // 共享内存按列主序读取到gpu数组中，冲突BDIMY次
}

__global__ void setRowReadColDynPad(int *out)
{
    extern __shared__ int tile[];                                           // 动态共享内存的申请需用extern关键字
    unsigned int row_idx = threadIdx.x + threadIdx.y * (blockDim.x + IPAD); // 按行展开
    unsigned int col_idx = threadIdx.y + threadIdx.x * (blockDim.y + IPAD); // 按列展开

    unsigned int g_idx = threadIdx.x + threadIdx.y * blockDim.x; // 获取编号索引

    tile[row_idx] = g_idx; // 按行主序设置，无冲突
    __syncthreads();       // 等待所有线程完成

    out[g_idx] = tile[col_idx]; // 共享内存按列主序读取到gpu数组中，由于进行了填充，无冲突
}

int main(int argc, char **argv)
{
    cudaSetDevice(DEVICENUM);
    int nElem = BDIMX * BDIMY;
    size_t nBytes = nElem * sizeof(int);

    // 设置用于接收方形共享内存的gpu数组
    int *device_array_res = nullptr;
    CHECK(cudaMalloc((int **)&device_array_res, nBytes));
    CHECK(cudaMemset(device_array_res, 0, nBytes));

    // 设置用于接收gpu结果数组的cpu指针
    int *gpu_res_to_cpu = (int *)malloc(nBytes);
    memset(gpu_res_to_cpu, 0, nBytes);

    dim3 gridSize(1, 1);
    dim3 blockSize(BDIMX, BDIMY);

    // 1. 设置动态共享内存，用一维模拟方形内存，按照行设置，按照列读取
    setRowReadColDyn<<<gridSize, blockSize, nElem * sizeof(int)>>>(device_array_res);
    CHECK(cudaMemcpy(gpu_res_to_cpu, device_array_res, nBytes, cudaMemcpyDeviceToHost));

    // 2. 设置动态共享内存，用一维模拟方形内存，按照行设置，按照列读取
    setRowReadColDynPad<<<gridSize, blockSize, (BDIMX + IPAD) * BDIMY * sizeof(int)>>>(device_array_res);
    CHECK(cudaMemcpy(gpu_res_to_cpu, device_array_res, nBytes, cudaMemcpyDeviceToHost));

    cudaFree(device_array_res);
    free(gpu_res_to_cpu);

    return 0;
}